
#include "kernel_operator.h"
// tensor num for each queue
constexpr int32_t BUFFER_NUM = 2;

// TYPE_X表示src的数据类型，支持四种情况，TYPE_Y表示index，数据类型为int64_t，
template<typename TYPE_X, typename TYPE_Y> class KernelAdd {
    using T = TYPE_X;
public:
    __aicore__ inline KernelAdd() {}
    // indicies表示为y，作为索引的情况
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR u, uint32_t smallCoreDataNum,
                                uint32_t bigCoreDataNum, uint32_t finalBigTileNum, 
                                uint32_t finalSmallTileNum, uint32_t tileDataNum, 
                                uint32_t smallTailDataNum, uint32_t bigTailDataNum, 
                                uint32_t tailBlockNum, int64_t axis, uint32_t indexLength,
                                uint32_t reduce, uint32_t indexLengthAligned,
                                uint32_t unitCount, uint32_t totalLength, uint32_t stepSize,
                                uint32_t dimSize, uint32_t firstDimIndex, uint32_t secondDimIndex, uint32_t thirdDimIndex,
                                uint32_t firstDimSelf, uint32_t secondDimSelf, uint32_t thirdDimSelf,
                                uint32_t firstDimUpdates, uint32_t secondDimUpdates, uint32_t thirdDimUpdates,
                                uint32_t updatesLength, uint32_t updatesLengthAligned,
                                uint32_t indexTileNum, uint32_t indexTailDataNum,
                                uint32_t updatesTileNum, uint32_t updatesTailDataNum) 
    {
        ASSERT(AscendC::GetBlockNum() != 0 && "block dim can not be zero!");
        uint32_t coreNum = AscendC::GetBlockIdx();
        uint32_t globalBufferIndex = bigCoreDataNum * AscendC::GetBlockIdx();
        this->tileDataNum = tileDataNum;
        this->indexLength = indexLength;
        this->indexLengthAligned = indexLengthAligned;
        this->updatesLength = updatesLength;
        this->updatesLengthAligned = updatesLengthAligned;
        // 表示待处理维度
        this->axis = axis;
        this->stepSize = stepSize;
        this->totalLength = totalLength;
        this->reduce = reduce;
        this->dimSize = dimSize;
        this->firstDimIndex = firstDimIndex;
        this->secondDimIndex = secondDimIndex;
        this->thirdDimIndex = thirdDimIndex;
        this->firstDimSelf = firstDimSelf;
        this->secondDimSelf = secondDimSelf;
        this->thirdDimSelf = thirdDimSelf;
        this->firstDimUpdates = firstDimUpdates;
        this->secondDimUpdates = secondDimUpdates;
        this->thirdDimUpdates = thirdDimUpdates;
        this->indexTileNum = indexTileNum;
        this->indexTailDataNum = indexTailDataNum;
        this->updatesTileNum = updatesTileNum;
        this->updatesTailDataNum = updatesTailDataNum;
        if (coreNum < tailBlockNum) { 
          this->coreDataNum = bigCoreDataNum;
          this->tileNum = finalBigTileNum;
          this->tailDataNum = bigTailDataNum;
        }
        else { 
          this->coreDataNum = smallCoreDataNum;
          this->tileNum = finalSmallTileNum;
          this->tailDataNum = smallTailDataNum;
          // globalBufferIndex -= (bigCoreDataNum - smallCoreDataNum) * (AscendC::GetBlockIdx() - tailBlockNum);
        }
        xGm.SetGlobalBuffer((__gm__ TYPE_X*)x, this->coreDataNum);
        // y和m的大小是一样的，但是和x的大小不一样
        yGm.SetGlobalBuffer((__gm__ TYPE_Y*)y, this->indexLengthAligned);
        uGm.SetGlobalBuffer((__gm__ TYPE_X*)u, this->updatesLengthAligned);

        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
        //注意此处给分配的内存空间要对齐，不要过大或过小
        if(this->indexLengthAligned < this->tileDataNum) {
          pipe.InitBuffer(inQueueY, BUFFER_NUM, this->indexLengthAligned * sizeof(TYPE_Y));
        } else {
          pipe.InitBuffer(inQueueY, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_Y));
        }
        if(this->updatesLengthAligned < this->tileDataNum) {
          pipe.InitBuffer(inQueueU, BUFFER_NUM, this->updatesLengthAligned * sizeof(TYPE_X));
        } else {
          pipe.InitBuffer(inQueueU, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
        }
        pipe.InitBuffer(outQueueZ, BUFFER_NUM, this->tileDataNum * sizeof(TYPE_X));
        
    }
    __aicore__ inline void Process()
    {
        int32_t loopCount = this->tileNum;
        this->processDataNum = this->tileDataNum;
        // 注意此处的顺序不要反了，外层的遍历是tile，这样设计可以确保每次CopyIn进行拷入拷出之后可以正确的完成一个tile的处理
        for (int32_t i = 0; i < loopCount; i++) {
          if (i == this->tileNum - 1) {
            this->processDataNum = this->tailDataNum;
            if constexpr (std::is_same_v<T, int32_t> || std::is_same_v<T, half>) {
              pipe.InitBuffer(tmp1, this->tileDataNum * sizeof(float));
              pipe.InitBuffer(tmp2, this->tileDataNum * sizeof(float));
            } else if constexpr (std::is_same_v<T, int8_t>) {
              pipe.InitBuffer(tmp1, this->tileDataNum * sizeof(half));
              pipe.InitBuffer(tmp2, this->tileDataNum * sizeof(half));
            }
          } else {
            if constexpr (std::is_same_v<T, int32_t> || std::is_same_v<T, half>) {
              pipe.InitBuffer(tmp1, this->tailDataNum * sizeof(float));
              pipe.InitBuffer(tmp2, this->tailDataNum * sizeof(float));
            } else if constexpr (std::is_same_v<T, int8_t>) {
              pipe.InitBuffer(tmp1, this->tailDataNum * sizeof(half));
              pipe.InitBuffer(tmp2, this->tailDataNum * sizeof(half));
            }
          }

          
          ProcessInner(i);
        }
    }

    __aicore__ inline void ProcessInner(int32_t i)
    {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
        
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
      AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.AllocTensor<TYPE_X>();
      AscendC::DataCopy(xLocal, xGm[progress * this->tileDataNum], this->processDataNum);
      inQueueX.EnQue(xLocal);
    }
    __aicore__ inline void Compute(int32_t progress)
    { 
      AscendC::LocalTensor<TYPE_X> xLocal = inQueueX.DeQue<TYPE_X>();
      AscendC::LocalTensor<TYPE_X> zLocal = outQueueZ.AllocTensor<TYPE_X>();
      if(this->dimSize > 3) {
        AscendC::DataCopy(zLocal, xLocal, this->processDataNum);
        return;
      }
      for(int32_t j = 0; j < this->indexTileNum; j++) {
        int32_t indexCopyLength = this->tileDataNum;
        AscendC::LocalTensor<TYPE_Y> yLocal = inQueueY.AllocTensor<TYPE_Y>();
        AscendC::LocalTensor<TYPE_X> uLocal = inQueueU.AllocTensor<TYPE_X>();
        if(j == this->indexTileNum - 1) {
          indexCopyLength = this->indexTailDataNum;
        }
        //读取y和uLocal
        AscendC::DataCopy(yLocal, yGm[j * this->tileDataNum], indexCopyLength);
        AscendC::DataCopy(uLocal, uGm[j * this->tileDataNum], indexCopyLength);
        // 内部处理

        if constexpr (std::is_same_v<T, uint8_t>) {
          // 使用half类型的计算
          auto zLocal_fp16 = tmp1.Get<half>();
          auto xLocal_fp16 = tmp2.Get<half>();
          AscendC::Cast(zLocal_fp16, xLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
          if(this->reduce == 0) {
            // uint8_t的Duplicate目前不支持
            AscendC::Duplicate<half>(xLocal_fp16, (half)0, this->processDataNum);
            for(int p = 0; p < indexCopyLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                zLocal_fp16.SetValue(localIndex, static_cast<half>((uLocal.GetValue(updatesLocalIndex)) + (xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Add(zLocal_fp16, zLocal_fp16, xLocal_fp16, this->processDataNum);
          }else {
            AscendC::Duplicate<half>(xLocal_fp16, (half)1, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);
              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                zLocal_fp16.SetValue(localIndex, static_cast<half>((uLocal.GetValue(updatesLocalIndex)) * (xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Mul(zLocal_fp16, zLocal_fp16, xLocal_fp16, this->processDataNum);
          }
          AscendC::Cast(zLocal, zLocal_fp16, AscendC::RoundMode::CAST_NONE, this->processDataNum);
        } else if constexpr (std::is_same_v<T, int32_t>) {
          AscendC::DataCopy(zLocal, xLocal, this->processDataNum);
          if(this->reduce == 0) {
            AscendC::Duplicate<int32_t>(xLocal, (half)0, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal.SetValue(localIndex, ((uLocal.GetValue(updatesLocalIndex)) + (xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Add(zLocal, zLocal, xLocal, this->processDataNum);
          }else {
            AscendC::Duplicate<int32_t>(xLocal, (half)1, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal.SetValue(localIndex, ((uLocal.GetValue(updatesLocalIndex)) * (xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Mul(zLocal, zLocal, xLocal, this->processDataNum);
          }
        } else if constexpr (std::is_same_v<T, half>) {
          // 进行矩阵变换
          // 通过实验已经证明了zLocal = xLocal不可以，会报错507046
          auto zLocal_fp32 = tmp1.Get<float>();
          auto xLocal_fp32 = tmp2.Get<float>();
          AscendC::Cast(zLocal_fp32, xLocal, AscendC::RoundMode::CAST_NONE, this->processDataNum);
          if(this->reduce == 0) {
            // 填充出现了错误
            AscendC::Duplicate<float>(xLocal_fp32, (float)0, this->processDataNum);
            // 考虑realIndex
            // 判断j是否在此loop的处理范围内
            for(int p = 0; p < indexCopyLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal_fp32.SetValue(localIndex, (static_cast<float>(uLocal.GetValue(updatesLocalIndex)) + static_cast<float>(xLocal_fp32.GetValue(localIndex))));
              }
            }
            AscendC::Add(zLocal_fp32, zLocal_fp32, xLocal_fp32, this->processDataNum);
          } else {
            AscendC::Duplicate<float>(xLocal_fp32, (float)1, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal_fp32.SetValue(localIndex, (static_cast<float>(uLocal.GetValue(updatesLocalIndex)) * static_cast<float>(xLocal_fp32.GetValue(localIndex))));
              }
            }
            AscendC::Mul(zLocal_fp32, zLocal_fp32, xLocal_fp32, this->processDataNum);
          }
          AscendC::Cast(zLocal, zLocal_fp32, AscendC::RoundMode::CAST_NONE, this->processDataNum);
          
        } else if constexpr (std::is_same_v<T, float>) {
          AscendC::DataCopy(zLocal, xLocal, this->processDataNum);
          if(this->reduce == 0) {
            AscendC::Duplicate<float>(xLocal, (half)0, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal.SetValue(localIndex, (static_cast<float>(uLocal.GetValue(updatesLocalIndex)) + static_cast<float>(xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Add(zLocal, zLocal, xLocal, this->processDataNum);
          }else {
            AscendC::Duplicate<float>(xLocal, (half)1, this->processDataNum);
            for(int p = 0; p < this->indexLength; p++) {
              // yLocal.GetValue(i)表示对应的位置
              int64_t i = p + j * this->tileDataNum;
              int64_t index = yLocal.GetValue(i);
              // 把一维的索引转成高维的索引
              int64_t l = i % this->thirdDimIndex;
              int64_t n = (i / this->thirdDimIndex) % this->secondDimIndex;
              int64_t m = (i / this->thirdDimIndex / this->secondDimIndex) % this->firstDimIndex;

              int64_t updatesIndex = m * this->secondDimUpdates * this->thirdDimUpdates + n * this->thirdDimUpdates + l;
              int64_t updatesLocalIndex = updatesIndex % (this->tileDataNum);

              // 根据m,n,l计算self和updates的位置
              int64_t realLocation = this->dimSize - this->axis;
              if(realLocation == 3) {
                m = index;
              } else if(realLocation == 2) {
                n = index;
              } else if(realLocation == 1) {
                l = index;
              }
              int64_t realIndex = m * this->secondDimSelf * this->thirdDimSelf + n * this->thirdDimSelf + l;

              if(realIndex >= progress * this->tileDataNum && realIndex < (progress + 1) * this->tileDataNum) {
                // 在当前tile处理
                int64_t localIndex = realIndex % (this->tileDataNum);
                xLocal.SetValue(localIndex, (static_cast<float>(uLocal.GetValue(updatesLocalIndex)) * static_cast<float>(xLocal.GetValue(localIndex))));
              }
            }
            AscendC::Mul(zLocal, zLocal, xLocal, this->processDataNum);
          }
        }

        //每次循环结束都会把zLocal赋值给xLocal
        AscendC::DataCopy(xLocal, zLocal, this->processDataNum);
        inQueueY.FreeTensor(yLocal);
        inQueueU.FreeTensor(uLocal);
      }
      outQueueZ.EnQue<TYPE_X>(zLocal);
      inQueueX.FreeTensor(xLocal);
    }

    __aicore__ inline void CopyOut(int32_t progress)
    {
      AscendC::LocalTensor<TYPE_X> zLocal = outQueueZ.DeQue<TYPE_X>();  
      AscendC::DataCopy(xGm[progress * this->tileDataNum], zLocal, this->processDataNum);
      outQueueZ.FreeTensor(zLocal);
    }

private:
    AscendC::TPipe pipe;
    AscendC::TQue<AscendC::QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueU;
    AscendC::TQue<AscendC::QuePosition::VECOUT, BUFFER_NUM> outQueueZ;
    AscendC::TBuf<AscendC::QuePosition::VECCALC> tmp1, tmp2;
    AscendC::GlobalTensor<TYPE_X> xGm;
    AscendC::GlobalTensor<TYPE_X> uGm;
    AscendC::GlobalTensor<TYPE_Y> yGm;
    uint32_t coreDataNum;
    uint32_t tileNum;
    uint32_t tileDataNum;
    uint32_t tailDataNum;
    uint32_t processDataNum;
    int64_t axis;
    uint32_t indexLength;
    uint32_t reduce;
    uint32_t indexLengthAligned;
    uint32_t updatesLength;
    uint32_t updatesLengthAligned;
    uint32_t stepSize;
    uint32_t totalLength;
    uint32_t dimSize, firstDimIndex, secondDimIndex, thirdDimIndex, firstDimSelf, secondDimSelf, thirdDimSelf, firstDimUpdates, secondDimUpdates, thirdDimUpdates;
    uint32_t indexTileNum, indexTailDataNum, updatesTileNum, updatesTailDataNum;
};

extern "C" __global__ __aicore__ void scatter_elements(GM_ADDR var, GM_ADDR indices, GM_ADDR updates, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    KernelAdd<DTYPE_VAR, DTYPE_INDICES> op;
    op.Init(var, indices, updates, tiling_data.smallCoreDataNum, 
            tiling_data.bigCoreDataNum, tiling_data.finalBigTileNum, 
            tiling_data.finalSmallTileNum, tiling_data.tileDataNum, 
            tiling_data.smallTailDataNum, tiling_data.bigTailDataNum, 
            tiling_data.tailBlockNum, tiling_data.axis, tiling_data.indexLength,
            tiling_data.reduce, tiling_data.indexLengthAligned,
            tiling_data.unitCount, tiling_data.totalLength, tiling_data.stepSize,
            tiling_data.dimSize, tiling_data.firstDimIndex, tiling_data.secondDimIndex, tiling_data.thirdDimIndex,
            tiling_data.firstDimSelf, tiling_data.secondDimSelf, tiling_data.thirdDimSelf,
            tiling_data.firstDimUpdates, tiling_data.secondDimUpdates, tiling_data.thirdDimUpdates,
            tiling_data.updatesLength, tiling_data.updatesLengthAligned,
            tiling_data.indexTileNum, tiling_data.indexTailDataNum,
            tiling_data.updatesTileNum, tiling_data.updatesTailDataNum);  
    op.Process();
}
